// phalanx scoped_ptr.hpp.cu header file

// Copyright (c) 2011 - 2011 Kohei Takahashi (Flast).
// Distributed under the MIT license. for more detail see COPYING.

#ifndef IG_PHALANX_TEST_DETAIL_SCOPED_PTR_HPP_CU_ONCE_
#define IG_PHALANX_TEST_DETAIL_SCOPED_PTR_HPP_CU_ONCE_

#include <phalanx/config.hpp.cu>

#include <cstddef>

#include <boost/mpl/map.hpp>
#include <boost/mpl/pair.hpp>
#include <boost/mpl/at.hpp>

#include <phalanx/detail/retypedef.hpp.cu>

#include <phalanx/utility/enable_if.hpp.cu>

#include <phalanx/type_traits/is_same.hpp.cu>

namespace phalanx
{

namespace test_detail
{

// Memory location
struct host;
struct device;
struct mapped;

typedef cudaError_t _free_func_t( void * );

template < typename T, typename _location >
struct _scoped_ptr_impl
{
    typedef T *              pointer;
    typedef T &              reference;
    typedef ::std::size_t    size_type;
    typedef ::std::ptrdiff_t difference_type;

protected:
    template < _free_func_t *_free_func >
    struct _caller
    {
        template < typename _T >
        PHALANX_CONFIG_ATTRIBUTE( host )
        static inline void
        _free( _T *ptr )
        { ptr ?  _free_func( ptr ) : 0; }
    };

    struct _nothing_to_do
    {
        template < typename _T >
        PHALANX_CONFIG_ATTRIBUTE( host )
        static inline void
        _free( _T * ) {}
    };

    typedef ::boost::mpl::map<
      ::boost::mpl::pair< host, _caller< cudaFreeHost > >,
      ::boost::mpl::pair< device, _caller< cudaFree > >,
      ::boost::mpl::pair< mapped, _nothing_to_do > > _free_func_map;

    size_type _size;
    pointer   _ptr;

    PHALANX_CONFIG_ATTRIBUTE( host )
    explicit
    _scoped_ptr_impl( size_type size, pointer ptr = 0 )
      : _size( size ), _ptr( ptr ) {}

    PHALANX_CONFIG_ATTRIBUTE( host )
    ~_scoped_ptr_impl( void )
    { reset(); }

    PHALANX_CONFIG_ATTRIBUTE( host )
    pointer &
    _get_ptr_instance( void )
    { return _ptr; }

public:
    PHALANX_CONFIG_ATTRIBUTE( host )
    void
    reset( pointer _new_ptr = 0 )
    {
        typedef typename ::boost::mpl::at< _free_func_map, _location >::type _caller_t;
        _caller_t::_free( _ptr );
        _ptr = _new_ptr;
    }

    PHALANX_CONFIG_ATTRIBUTE( host )
    pointer
    get( void ) const
    { return _ptr; }

    PHALANX_CONFIG_ATTRIBUTE( host )
    size_type
    instance_size( void ) const
    { return sizeof( T ) * size(); }

    PHALANX_CONFIG_ATTRIBUTE( host )
    size_type
    size( void ) const
    { return _size; }

    PHALANX_CONFIG_ATTRIBUTE( host )
    reference
    operator*( void ) const
    { return *get(); }

    PHALANX_CONFIG_ATTRIBUTE( host )
    reference
    operator[]( difference_type _x ) const
    { return get()[ _x ]; }
};

template < typename T, typename _location >
struct scoped_ptr
  : public _scoped_ptr_impl< T, _location >
{
    typedef _scoped_ptr_impl< T, _location > _base_t;
    PHALANX_PP_RETYPEDEF( _base_t, pointer );
    PHALANX_PP_RETYPEDEF( _base_t, reference );
    PHALANX_PP_RETYPEDEF( _base_t, size_type );
    PHALANX_PP_RETYPEDEF( _base_t, difference_type );

private:
    template < typename _loc >
    PHALANX_CONFIG_ATTRIBUTE( host )
    typename enable_if< is_same< _loc, host > >::type
    _allocator( unsigned int flags )
    {
        // FIXME: should check return code with assertion.
        cudaHostAlloc( &this->_get_ptr_instance(), this->instance_size(), flags );
    }

    template < typename _loc >
    PHALANX_CONFIG_ATTRIBUTE( host )
    typename enable_if< is_same< _loc, device > >::type
    _allocator( void )
    {
        // FIXME: should check return code with assertion.
        cudaMalloc( &this->_get_ptr_instance(), this->instance_size() );
    }

    template < typename _loc >
    PHALANX_CONFIG_ATTRIBUTE( host )
    typename enable_if< is_same< _loc, mapped > >::type
    _allocator( pointer _host_ptr )
    {
        // FIXME: should check return code with assertion.
        cudaHostGetDevicePointer( &this->_get_ptr_instance(), _host_ptr, 0 );
    }

public:
    PHALANX_CONFIG_ATTRIBUTE( host )
    scoped_ptr( void )
      : _base_t( 0 ) {}

    template < typename _T >
    PHALANX_CONFIG_ATTRIBUTE( host )
    scoped_ptr( _T size, unsigned int flags )
      : _base_t( size )
    { _allocator< _location >( flags ); }

    template < typename _T >
    PHALANX_CONFIG_ATTRIBUTE( host )
    explicit
    scoped_ptr( _T size )
      : _base_t( size )
    { _allocator< _location >(); }

    PHALANX_CONFIG_ATTRIBUTE( host )
    explicit
    scoped_ptr( const scoped_ptr< T, host > &_host_ptr )
      : _base_t( _host_ptr.size() )
    { _allocator< _location >( _host_ptr.get() ); }
};

} // namespace test_detail

} // namespace phalanx

#endif // IG_PHALANX_TEST_DETAIL_SCOPED_PTR_HPP_CU_ONCE_

